#include "sat.h"
#include "scan.h"
#include "transpose.h"
#include "hr_time.h"

//#define DEBUG 

#ifdef DEBUG
void print_arr(SAT_test_t * t, float4 *arr);
#endif

enum Run {Scan, Transpose};

void split_time(SAT_test_t *t, Run r, CStopWatch * s) {
		cudaThreadSynchronize();
		s->stopTimer();
		if (r == Scan)
			t->scan_time += s->getElapsedTime();
		else if (r == Transpose)
			t->transpose_time += s->getElapsedTime();
}

double launch_SAT(SAT_test_t *t)
{
		CStopWatch tim;

#ifdef CUDA_TIME		
		cudaEvent_t start, stop;

		cudaEventCreate(&start);
		cudaEventCreate(&stop);
		cudaEventRecord( start, 0 );
#endif

#ifdef DEBUG
		float4 * tempa = (float4 *) malloc (t->num * t->num * sizeof(float4));
		cutilSafeCall(cudaMemcpy(tempa, t->dev_src, t->num*t->num*sizeof(float4), cudaMemcpyDeviceToHost));
		//for (int i = 0; i < 16; i++)
		//	printf("%3.0f, ", (*(tempa + 4351 + i)).x);
		printf("\n");
#endif

		printf("Copying input matrix to GPU...\n");
		cutilSafeCall(cudaMemcpy(t->dev_src, t->hst_src, t->num * t->num * sizeof(float4), cudaMemcpyHostToDevice));

		CStopWatch s1;
		s1.startTimer();
		float4 * temp;
		tim.startTimer();
		temp = sat_scan<float4>(t, t->num, t->num);
		if (temp == t->dev_src) {
			t->dev_src = t->dev_dst;
			t->dev_dst = temp;
		}
		if (t->split) {
			split_time(t, Scan, &s1);
		}

#ifdef DEBUG
		cutilSafeCall(cudaMemcpy(tempa, t->dev_dst, t->num*t->num*sizeof(float4), cudaMemcpyDeviceToHost));
		print_arr(t, tempa);
		printf("\n");
#endif

		CStopWatch s2;
		s2.startTimer();
		launch_naive_transpose(t->dev_dst, t->dev_src, t->num, t->num);
		if (t->split) {
			split_time(t, Transpose, &s2);
		}

#ifdef DEBUG
		cutilSafeCall(cudaMemcpy(tempa, t->dev_src, t->num*t->num*sizeof(float4), cudaMemcpyDeviceToHost));
		print_arr(t, tempa);
		printf("\n");
#endif

		CStopWatch s3;
		s3.startTimer();
		temp = sat_scan<float4>(t, t->num, t->num);
		if (temp == t->dev_src) {
			t->dev_src = t->dev_dst;
			t->dev_dst = temp;
		}
		if (t->split) {
			split_time(t, Scan, &s3);
		}

#ifdef DEBUG
		cutilSafeCall(cudaMemcpy(tempa, t->dev_dst, t->num*t->num*sizeof(float4), cudaMemcpyDeviceToHost));
		print_arr(t, tempa);
		printf("\n");
#endif

		CStopWatch s4;
		s4.startTimer();
		launch_naive_transpose(t->dev_dst, t->dev_src, t->num, t->num);
		temp = t->dev_dst;
		t->dev_dst = t->dev_src;
		t->dev_src = temp;
		if (t->split) {
			split_time(t, Transpose, &s4);
		}

		cutilSafeCall(cudaThreadSynchronize());
		tim.stopTimer();

#ifdef CUDA_TIME
		cudaEventRecord( stop, 0 );
		cudaEventSynchronize( stop );
		cudaEventElapsedTime( &time, start, stop );
		cudaEventDestroy( start );
		cudaEventDestroy( stop );
#endif

#ifdef DEBUG
		cutilSafeCall(cudaMemcpy(tempa, t->dev_dst, t->num*t->num*sizeof(float4), cudaMemcpyDeviceToHost));
		print_arr(t, tempa);
		printf("\n");
		free(tempa);
#endif

		return tim.getElapsedTime();
}
